

#define DIV255 0.00392156862745098
#define DIV25512  0.000326797385620915


#if 0
kernel void dict_bgr2z(global const float *dict,global const char* buff, global float* dst)  
{  

    int x = (int)get_global_id(0);  
    int y = (int)get_global_id(1);  

    unsigned char b_tmp = buff[(x*800 + y)*4];
    unsigned char g_tmp = buff[(x*800 + y)*4 + 1];
    unsigned char r_tmp = buff[(x*800 + y)*4 + 2];
    float b = (float)b_tmp ;
    float g = (float)g_tmp ;
    float r = (float)r_tmp ;
    
//    printf("b:%f g:%f r:%f\n",b,g,r);
	if (r <= 10.0f)
	{
		r = r / (255.0f*12.0f) ;
		g = g / (255.0f) ;
		b = b / (255.0f) ;
	}
	else
	{
		r = dict[(int)r] ;
		
		if (g <= 10.0f)
		{
			g = g / (255.0f * 12.0f) ;
			b = b / (255.0f) ;
		}
		else
		{
			g = dict[(int)g] ;
			if (b <= 10.0f)
			{
				b = b / (255.0f * 12.0f);
			}
			else
			{
				b = dict[(int)b] ;
			}
		}
	}

	float Z = 0.019334f * r + 0.119193f * g + 0.950527f * b;
	
	//dot( (float3)(0.019334f,  0.119193f, 0.950527f), (float3)(r,g, b));
	
//	printf("b:%f g:%f r:%f Z:%f\n",b,g,r,Z);
	dst[x+y*1280] = Z ;
}  

#else

//kernel void dict_bgr2z(__constant float *dict,global __read_only uchar4* buff, global __write_only float* dst)  
//{  
//    uint idxIo = (get_global_id(0)*800 + get_global_id(1));
//    uint idxOut = (get_global_id(0) + get_global_id(1)*1280);
//    
//    float3 bgr;
//
//    bgr = convert_float3(buff[idxIo].s012);
//    
//    //bgr.s0 = ((bgr.s0+14)/269)*((bgr.s0+14)/269);
//	//bgr.s1 = ((bgr.s1+14)/269)*((bgr.s1+14)/269);
//	//bgr.s2 = ((bgr.s2+14)/269)*((bgr.s2+14)/269);
//#define COEF_1D269  0.003717472118959f
//#define COEF_14D269 0.052044609665428f
//    
//    bgr = bgr*COEF_1D269  + COEF_14D269;
//
//	dst[idxOut] = dot( (float3)(0.013929122f, 0.097097002f, 0.71418547f), bgr*bgr);
//	
//}  


// one work item computes 4*1 pixels
kernel void dict_bgr2z(__constant float *dict,
		               __global __read_only  uchar4* buff, 
					   __global __write_only float4* dst)  
{  
	uint ix = get_global_id(0);
	uint iy = get_global_id(1);
	
#define WIDTH   (960)
	
	uint idxI = ix * 3 + iy * WIDTH;
	float4  bgrb = convert_float4(buff[idxI]);
	float4  grbg = convert_float4(buff[++idxI]);
	float4  rbgr = convert_float4(buff[++idxI]);

#define COEF_1D269  0.003717472118959f
#define COEF_14D269 0.052044609665428f
    
	bgrb = bgrb*COEF_1D269  + COEF_14D269;
	grbg = grbg*COEF_1D269  + COEF_14D269;
	rbgr = rbgr*COEF_1D269  + COEF_14D269;
	
	bgrb = bgrb*bgrb;
	grbg = grbg*grbg;
	rbgr = rbgr*rbgr;
	
	
	float4 out;
	
	const float3 dot_coef = (float3)(0.013929122f, 0.097097002f, 0.71418547f);
	
	out.s0 = dot(dot_coef , bgrb.s012);
	out.s1 = dot(dot_coef , (float3)(bgrb.s3, grbg.s01) );
	out.s2 = dot(dot_coef , (float3)(grbg.s23, rbgr.s0) );
	out.s3 = dot(dot_coef , (rbgr.s123) );
	

	dst[ix + iy * 320] = out ;
	
}


//// one work item computes 1*1 pixels
//kernel void dict_bgr2z(__constant float *dict,
//		               __global __read_only  uchar* buff, 
//					   __global __write_only float* dst)  
//{  
//	uint ix = get_global_id(0);
//	uint iy = get_global_id(1);
//	
//#define WIDTH   (1280*3)
//	
//	uint idxI = ix * 3 + iy * WIDTH;
//	float3  bgr = convert_float3(vload3(0, buff+idxI));
//
//
//#define COEF_1D269  0.003717472118959f
//#define COEF_14D269 0.052044609665428f
//    
//	bgr = bgr*COEF_1D269  + COEF_14D269;
//
//	dst[ix + iy * 1280] = dot((float3)(0.013929122f, 0.097097002f, 0.71418547f),  bgr*bgr) ;
//}


//// one work item computes 1*1 pixels
//kernel void dict_bgr2z(__constant float *dict,
//		               __global __read_only  uchar* buff, 
//					   __global __write_only float* dst)  
//{  
//	uint ix = get_global_id(0);
//	uint iy = get_global_id(1);
//	
//#define WIDTH   (1280*3)
//	
//	uint idxI = ix * 3 *800 + iy * 3;
//	float3  bgr = convert_float3(vload3(0, buff+idxI));
//
//
//#define COEF_1D269  0.003717472118959f
//#define COEF_14D269 0.052044609665428f
//    
//	bgr = bgr*COEF_1D269  + COEF_14D269;
//
//	dst[ix + iy * 1280] = dot((float3)(0.013929122f, 0.097097002f, 0.71418547f),  bgr*bgr) ;
//}

#endif
